﻿/*
    Alunos:
        Jorge Luis Stefanelli Bonafé	5889794
        Marcelo Koti Kamada             7277670
*/

#include <stdio.h>
#include <stdlib.h>
#include <cmath>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

// Processa a imagem
// Nas bordas, alguns pixels adjacentes valem -1 (fora da imagem).
// Esses pixels não são levados em conta no cálculo da média.
// nDiv é o número total de pixels válidos, que é também
// o valor da divisão na média aritmetica.
__global__ void kernel(int *inR, int *inG, int *inB, int nRow, int nCol) {
    int tc = threadIdx.y;           // coluna da matriz de threadas
    int tl = threadIdx.x;           // linha da matriz

    int bc = blockIdx.y;            // coluna da matriz de blocos
    int bl = blockIdx.x;            // linha da matriz

    int pc = tc + bc*blockDim.y;    // calculo da coluna do pixel
	int pl = tl*nCol + bl*16*nCol;  // calculo da linha do pixel

    if(pc < nCol && pl < nRow*nCol) {
        int rAdj;                    // valor da linha do pixel adjacente ao processado
        int cAdj;                    // valor da coluna do pixel adjacente ao processado
        int totalR = 0;              // Valor total da soma dos 5x5 pixels para cada componente de cor
        int totalG = 0;              // componente verde
        int totalB = 0;              // componente azul
        int nDiv = 0;                // numero de pixeis que estao sendo somados

        for(rAdj = - 2; rAdj <= 2; rAdj++) {
            int lAtual = pl+(rAdj*nCol);
            if(lAtual<0 || lAtual>=nRow*nCol) continue; // impede que os blocos que tem pedacos fora da imagem facam calculos

            for(cAdj = -2; cAdj <= 2; cAdj++) {
                int cAtual = pc+cAdj;
                if(cAtual < 0 || cAtual > nCol) continue; // mesma checagem que a anterior

                totalR += inR[lAtual+cAtual];
                totalG += inG[lAtual+cAtual];
                totalB += inB[lAtual+cAtual];
                
                nDiv++;
            }
        }
        __syncthreads();                    // que todas threads terminem para atualizar a matriz
        inR[pl+pc] = totalR/nDiv;           // atualiza a matriz de cada cor
        inG[pl+pc] = totalG/nDiv;
        inB[pl+pc] = totalB/nDiv;
    }
}

int main(int argc, char *argv[]) {

	// ----- leitura do cabecalho do in, escrita do cabecalho do out -----
	int maxColor;       // nao e' usado
	int nRow, nCol;     // numero de linhas e numero de colunas respectivamente
	char sIn[10];       // tipo do arquivo P3 ou P6

    scanf("%s",sIn);            // Lê P3 ou P6
    scanf("%d%d",&nCol,&nRow);  // Lê Nº de colunas e de linhas
    scanf("%d",&maxColor);      // Lê valor máximo de cor (não usado)

	// ----- alocacao das matrizes -----

	int *R = (int*) malloc(sizeof(int) * nRow * nCol);
	int *G = (int*) malloc(sizeof(int) * nRow * nCol);
	int *B = (int*) malloc(sizeof(int) * nRow * nCol);

	int *dev_R, *dev_G, *dev_B;
	cudaMalloc( (void**) &dev_R, sizeof(int) * nRow * nCol);
	cudaMalloc( (void**) &dev_G, sizeof(int) * nRow * nCol);
	cudaMalloc( (void**) &dev_B, sizeof(int) * nRow * nCol);

	// ----- Leitura da imagem de entrada -----
	int i, j;
	for(i = 0; i < nRow; i++) {
		for(j = 0; j < nCol; j++) {
            scanf("%d%d%d", &R[nCol*i + j], &G[nCol*i + j], &B[nCol*i + j]);
		}
	}

	// ----- copia as matrizes para a memoria da GPU -----
	cudaMemcpy(dev_R, R, sizeof(int) * nRow * nCol, cudaMemcpyHostToDevice);
	cudaMemcpy(dev_G, G, sizeof(int) * nRow * nCol, cudaMemcpyHostToDevice);
	cudaMemcpy(dev_B, B, sizeof(int) * nRow * nCol, cudaMemcpyHostToDevice);

	int hBlocks = ceil(((float)nCol / 16.0));   // numero de blocos na horizontal no grid (colunas)
	int vBlocks = ceil(((float)nRow / 16.0));   // numero de blocos na vertical no grid   (linhas)

	dim3 threads(16, 16);                       // cada bloco tem uma matriz de 16x16 threads
	dim3 blocks(vBlocks, hBlocks);              // o grid tem uma matriz de vBlocks x hBlocks blocos

	//// ----- faz o processamento da imagem -----
    kernel<<<blocks, threads>>>(dev_R, dev_G, dev_B, nRow, nCol);

	//// ----- copia de volta o resultado da GPU para a memoria RAM -----
	cudaMemcpy(R, dev_R, sizeof(int) * nRow * nCol, cudaMemcpyDeviceToHost);
	cudaMemcpy(G, dev_G, sizeof(int) * nRow * nCol, cudaMemcpyDeviceToHost);
	cudaMemcpy(B, dev_B, sizeof(int) * nRow * nCol, cudaMemcpyDeviceToHost);

	// ----- Imprime o arquivo de saida -----
    printf("P3\n");
    printf("%d %d\n%d\n",nCol,nRow,maxColor);
	for(i = 0; i < nRow; i++) {
		for(j = 0; j < nCol; j++)
			printf("%d %d %d ", R[nCol*i + j], G[nCol*i + j], B[nCol*i + j]);
		printf("\n");
	}

	// ----- desalocacao de memoria -----
	free(R);
	free(G);
	free(B);

	cudaFree(dev_R);
	cudaFree(dev_G);
	cudaFree(dev_B);

	return 0;
}
